// This file is part of Eigen, a lightweight C++ template library
// for linear algebra.
//
// Copyright (C) 2014 Navdeep Jaitly <ndjaitly@google.com>
//                    Benoit Steiner <benoit.steiner.goog@gmail.com>
//
// This Source Code Form is subject to the terms of the Mozilla
// Public License v. 2.0. If a copy of the MPL was not distributed
// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.

#ifndef EIGEN_CXX11_TENSOR_TENSOR_REVERSE_H
#define EIGEN_CXX11_TENSOR_TENSOR_REVERSE_H
namespace Eigen {

/** \class TensorReverse
 * \ingroup CXX11_Tensor_Module
 *
 * \brief Tensor reverse elements class.
 *
 */
namespace internal {
template<typename ReverseDimensions, typename XprType>
struct traits<TensorReverseOp<ReverseDimensions, XprType>> : public traits<XprType>
{
	typedef typename XprType::Scalar Scalar;
	typedef traits<XprType> XprTraits;
	typedef typename XprTraits::StorageKind StorageKind;
	typedef typename XprTraits::Index Index;
	typedef typename XprType::Nested Nested;
	typedef typename remove_reference<Nested>::type _Nested;
	static const int NumDimensions = XprTraits::NumDimensions;
	static const int Layout = XprTraits::Layout;
	typedef typename XprTraits::PointerType PointerType;
};

template<typename ReverseDimensions, typename XprType>
struct eval<TensorReverseOp<ReverseDimensions, XprType>, Eigen::Dense>
{
	typedef const TensorReverseOp<ReverseDimensions, XprType>& type;
};

template<typename ReverseDimensions, typename XprType>
struct nested<TensorReverseOp<ReverseDimensions, XprType>,
			  1,
			  typename eval<TensorReverseOp<ReverseDimensions, XprType>>::type>
{
	typedef TensorReverseOp<ReverseDimensions, XprType> type;
};

} // end namespace internal

template<typename ReverseDimensions, typename XprType>
class TensorReverseOp : public TensorBase<TensorReverseOp<ReverseDimensions, XprType>, WriteAccessors>
{
  public:
	typedef TensorBase<TensorReverseOp<ReverseDimensions, XprType>, WriteAccessors> Base;
	typedef typename Eigen::internal::traits<TensorReverseOp>::Scalar Scalar;
	typedef typename Eigen::NumTraits<Scalar>::Real RealScalar;
	typedef typename XprType::CoeffReturnType CoeffReturnType;
	typedef typename Eigen::internal::nested<TensorReverseOp>::type Nested;
	typedef typename Eigen::internal::traits<TensorReverseOp>::StorageKind StorageKind;
	typedef typename Eigen::internal::traits<TensorReverseOp>::Index Index;

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorReverseOp(const XprType& expr, const ReverseDimensions& reverse_dims)
		: m_xpr(expr)
		, m_reverse_dims(reverse_dims)
	{
	}

	EIGEN_DEVICE_FUNC
	const ReverseDimensions& reverse() const { return m_reverse_dims; }

	EIGEN_DEVICE_FUNC
	const typename internal::remove_all<typename XprType::Nested>::type& expression() const { return m_xpr; }

	EIGEN_TENSOR_INHERIT_ASSIGNMENT_OPERATORS(TensorReverseOp)

  protected:
	typename XprType::Nested m_xpr;
	const ReverseDimensions m_reverse_dims;
};

// Eval as rvalue
template<typename ReverseDimensions, typename ArgType, typename Device>
struct TensorEvaluator<const TensorReverseOp<ReverseDimensions, ArgType>, Device>
{
	typedef TensorReverseOp<ReverseDimensions, ArgType> XprType;
	typedef typename XprType::Index Index;
	static const int NumDims = internal::array_size<ReverseDimensions>::value;
	typedef DSizes<Index, NumDims> Dimensions;
	typedef typename XprType::Scalar Scalar;
	typedef typename XprType::CoeffReturnType CoeffReturnType;
	typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
	static const int PacketSize = PacketType<CoeffReturnType, Device>::size;
	typedef StorageMemory<CoeffReturnType, Device> Storage;
	typedef typename Storage::Type EvaluatorPointerType;

	enum
	{
		IsAligned = false,
		PacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess,
		BlockAccess = NumDims > 0,
		PreferBlockAccess = true,
		Layout = TensorEvaluator<ArgType, Device>::Layout,
		CoordAccess = false, // to be implemented
		RawAccess = false
	};

	typedef internal::TensorIntDivisor<Index> IndexDivisor;

	//===- Tensor block evaluation strategy (see TensorBlock.h) -------------===//
	typedef internal::TensorBlockDescriptor<NumDims, Index> TensorBlockDesc;
	typedef internal::TensorBlockScratchAllocator<Device> TensorBlockScratch;

	typedef typename TensorEvaluator<const ArgType, Device>::TensorBlock ArgTensorBlock;

	typedef typename internal::TensorMaterializedBlock<CoeffReturnType, NumDims, Layout, Index> TensorBlock;
	//===--------------------------------------------------------------------===//

	EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
		: m_impl(op.expression(), device)
		, m_reverse(op.reverse())
		, m_device(device)
	{
		// Reversing a scalar isn't supported yet. It would be a no-op anyway.
		EIGEN_STATIC_ASSERT((NumDims > 0), YOU_MADE_A_PROGRAMMING_MISTAKE);

		// Compute strides
		m_dimensions = m_impl.dimensions();
		if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
			m_strides[0] = 1;
			for (int i = 1; i < NumDims; ++i) {
				m_strides[i] = m_strides[i - 1] * m_dimensions[i - 1];
				if (m_strides[i] > 0)
					m_fastStrides[i] = IndexDivisor(m_strides[i]);
			}
		} else {
			m_strides[NumDims - 1] = 1;
			for (int i = NumDims - 2; i >= 0; --i) {
				m_strides[i] = m_strides[i + 1] * m_dimensions[i + 1];
				if (m_strides[i] > 0)
					m_fastStrides[i] = IndexDivisor(m_strides[i]);
			}
		}
	}

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Dimensions& dimensions() const { return m_dimensions; }

	EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(EvaluatorPointerType)
	{
		m_impl.evalSubExprsIfNeeded(NULL);
		return true;
	}

#ifdef EIGEN_USE_THREADS
	template<typename EvalSubExprsCallback>
	EIGEN_STRONG_INLINE void evalSubExprsIfNeededAsync(EvaluatorPointerType, EvalSubExprsCallback done)
	{
		m_impl.evalSubExprsIfNeededAsync(nullptr, [done](bool) { done(true); });
	}
#endif // EIGEN_USE_THREADS

	EIGEN_STRONG_INLINE void cleanup() { m_impl.cleanup(); }

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index reverseIndex(Index index) const
	{
		eigen_assert(index < dimensions().TotalSize());
		Index inputIndex = 0;
		if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
			EIGEN_UNROLL_LOOP
			for (int i = NumDims - 1; i > 0; --i) {
				Index idx = index / m_fastStrides[i];
				index -= idx * m_strides[i];
				if (m_reverse[i]) {
					idx = m_dimensions[i] - idx - 1;
				}
				inputIndex += idx * m_strides[i];
			}
			if (m_reverse[0]) {
				inputIndex += (m_dimensions[0] - index - 1);
			} else {
				inputIndex += index;
			}
		} else {
			EIGEN_UNROLL_LOOP
			for (int i = 0; i < NumDims - 1; ++i) {
				Index idx = index / m_fastStrides[i];
				index -= idx * m_strides[i];
				if (m_reverse[i]) {
					idx = m_dimensions[i] - idx - 1;
				}
				inputIndex += idx * m_strides[i];
			}
			if (m_reverse[NumDims - 1]) {
				inputIndex += (m_dimensions[NumDims - 1] - index - 1);
			} else {
				inputIndex += index;
			}
		}
		return inputIndex;
	}

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index) const
	{
		return m_impl.coeff(reverseIndex(index));
	}

	template<int LoadMode>
	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType packet(Index index) const
	{
		EIGEN_STATIC_ASSERT((PacketSize > 1), YOU_MADE_A_PROGRAMMING_MISTAKE)
		eigen_assert(index + PacketSize - 1 < dimensions().TotalSize());

		// TODO(ndjaitly): write a better packing routine that uses
		// local structure.
		EIGEN_ALIGN_MAX typename internal::remove_const<CoeffReturnType>::type values[PacketSize];
		EIGEN_UNROLL_LOOP
		for (int i = 0; i < PacketSize; ++i) {
			values[i] = coeff(index + i);
		}
		PacketReturnType rslt = internal::pload<PacketReturnType>(values);
		return rslt;
	}

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE internal::TensorBlockResourceRequirements getResourceRequirements() const
	{
		const size_t target_size = m_device.lastLevelCacheSize();
		// Block evaluation reads underlying memory in reverse order, and default
		// cost model does not properly catch this in bytes stored/loaded.
		return internal::TensorBlockResourceRequirements::skewed<Scalar>(target_size).addCostPerCoeff({ 0, 0, 24 });
	}

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorBlock block(TensorBlockDesc& desc,
															TensorBlockScratch& scratch,
															bool /*root_of_expr_ast*/ = false) const
	{
		// TODO(ezhulenev): If underlying tensor expression supports and prefers
		// block evaluation we must use it. Currently we use coeff and packet
		// access into the underlying tensor expression.
		// static const bool useBlockAccessForArgType =
		//     TensorEvaluator<ArgType, Device>::BlockAccess &&
		//     TensorEvaluator<ArgType, Device>::PreferBlockAccess;

		static const bool isColMajor = static_cast<int>(Layout) == static_cast<int>(ColMajor);

		static const Index inner_dim_idx = isColMajor ? 0 : NumDims - 1;
		const bool inner_dim_reversed = m_reverse[inner_dim_idx];

		// Offset in the output block.
		Index block_offset = 0;

		// Offset in the input Tensor.
		Index input_offset = reverseIndex(desc.offset());

		// Initialize output block iterator state. Dimension in this array are
		// always in inner_most -> outer_most order (col major layout).
		array<BlockIteratorState, NumDims> it;
		for (int i = 0; i < NumDims; ++i) {
			const int dim = isColMajor ? i : NumDims - 1 - i;
			it[i].size = desc.dimension(dim);
			it[i].count = 0;
			it[i].reverse = m_reverse[dim];

			it[i].block_stride = i == 0 ? 1 : (it[i - 1].size * it[i - 1].block_stride);
			it[i].block_span = it[i].block_stride * (it[i].size - 1);

			it[i].input_stride = m_strides[dim];
			it[i].input_span = it[i].input_stride * (it[i].size - 1);

			if (it[i].reverse) {
				it[i].input_stride = -1 * it[i].input_stride;
				it[i].input_span = -1 * it[i].input_span;
			}
		}

		// If multiple inner dimensions have the same reverse flag, check if we can
		// merge them into a single virtual inner dimension.
		int effective_inner_dim = 0;
		for (int i = 1; i < NumDims; ++i) {
			if (it[i].reverse != it[effective_inner_dim].reverse)
				break;
			if (it[i].block_stride != it[effective_inner_dim].size)
				break;
			if (it[i].block_stride != numext::abs(it[i].input_stride))
				break;

			it[i].size = it[effective_inner_dim].size * it[i].size;

			it[i].block_stride = 1;
			it[i].input_stride = (inner_dim_reversed ? -1 : 1);

			it[i].block_span = it[i].block_stride * (it[i].size - 1);
			it[i].input_span = it[i].input_stride * (it[i].size - 1);

			effective_inner_dim = i;
		}

		eigen_assert(it[effective_inner_dim].block_stride == 1);
		eigen_assert(it[effective_inner_dim].input_stride == (inner_dim_reversed ? -1 : 1));

		const Index inner_dim_size = it[effective_inner_dim].size;

		// Prepare storage for the materialized reverse result.
		const typename TensorBlock::Storage block_storage = TensorBlock::prepareStorage(desc, scratch);
		CoeffReturnType* block_buffer = block_storage.data();

		while (it[NumDims - 1].count < it[NumDims - 1].size) {
			// Copy inner-most dimension data from reversed location in input.
			Index dst = block_offset;
			Index src = input_offset;

			// NOTE(ezhulenev): Adding vectorized path with internal::preverse showed
			// worse results in benchmarks than a simple coefficient loop.
			if (inner_dim_reversed) {
				for (Index i = 0; i < inner_dim_size; ++i) {
					block_buffer[dst] = m_impl.coeff(src);
					++dst;
					--src;
				}
			} else {
				for (Index i = 0; i < inner_dim_size; ++i) {
					block_buffer[dst] = m_impl.coeff(src);
					++dst;
					++src;
				}
			}

			// For the 1d tensor we need to generate only one inner-most dimension.
			if ((NumDims - effective_inner_dim) == 1)
				break;

			// Update offset.
			for (Index i = effective_inner_dim + 1; i < NumDims; ++i) {
				if (++it[i].count < it[i].size) {
					block_offset += it[i].block_stride;
					input_offset += it[i].input_stride;
					break;
				}
				if (i != NumDims - 1)
					it[i].count = 0;
				block_offset -= it[i].block_span;
				input_offset -= it[i].input_span;
			}
		}

		return block_storage.AsTensorMaterializedBlock();
	}

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(bool vectorized) const
	{
		double compute_cost = NumDims * (2 * TensorOpCost::AddCost<Index>() + 2 * TensorOpCost::MulCost<Index>() +
										 TensorOpCost::DivCost<Index>());
		for (int i = 0; i < NumDims; ++i) {
			if (m_reverse[i]) {
				compute_cost += 2 * TensorOpCost::AddCost<Index>();
			}
		}
		return m_impl.costPerCoeff(vectorized) + TensorOpCost(0, 0, compute_cost, false /* vectorized */, PacketSize);
	}

	EIGEN_DEVICE_FUNC typename Storage::Type data() const { return NULL; }

#ifdef EIGEN_USE_SYCL
	// binding placeholder accessors to a command group handler for SYCL
	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void bind(cl::sycl::handler& cgh) const { m_impl.bind(cgh); }
#endif

  protected:
	Dimensions m_dimensions;
	array<Index, NumDims> m_strides;
	array<IndexDivisor, NumDims> m_fastStrides;
	TensorEvaluator<ArgType, Device> m_impl;
	ReverseDimensions m_reverse;
	const Device EIGEN_DEVICE_REF m_device;

  private:
	struct BlockIteratorState
	{
		BlockIteratorState()
			: size(0)
			, count(0)
			, reverse(false)
			, block_stride(0)
			, block_span(0)
			, input_stride(0)
			, input_span(0)
		{
		}

		Index size;
		Index count;
		bool reverse;
		Index block_stride;
		Index block_span;
		Index input_stride;
		Index input_span;
	};
};

// Eval as lvalue

template<typename ReverseDimensions, typename ArgType, typename Device>
struct TensorEvaluator<TensorReverseOp<ReverseDimensions, ArgType>, Device>
	: public TensorEvaluator<const TensorReverseOp<ReverseDimensions, ArgType>, Device>
{
	typedef TensorEvaluator<const TensorReverseOp<ReverseDimensions, ArgType>, Device> Base;
	typedef TensorReverseOp<ReverseDimensions, ArgType> XprType;
	typedef typename XprType::Index Index;
	static const int NumDims = internal::array_size<ReverseDimensions>::value;
	typedef DSizes<Index, NumDims> Dimensions;

	enum
	{
		IsAligned = false,
		PacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess,
		BlockAccess = false,
		PreferBlockAccess = false,
		Layout = TensorEvaluator<ArgType, Device>::Layout,
		CoordAccess = false, // to be implemented
		RawAccess = false
	};
	EIGEN_STRONG_INLINE TensorEvaluator(const XprType& op, const Device& device)
		: Base(op, device)
	{
	}

	typedef typename XprType::Scalar Scalar;
	typedef typename XprType::CoeffReturnType CoeffReturnType;
	typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
	static const int PacketSize = PacketType<CoeffReturnType, Device>::size;

	//===- Tensor block evaluation strategy (see TensorBlock.h) -------------===//
	typedef internal::TensorBlockNotImplemented TensorBlock;
	//===--------------------------------------------------------------------===//

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Dimensions& dimensions() const { return this->m_dimensions; }

	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar& coeffRef(Index index)
	{
		return this->m_impl.coeffRef(this->reverseIndex(index));
	}

	template<int StoreMode>
	EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void writePacket(Index index, const PacketReturnType& x)
	{
		EIGEN_STATIC_ASSERT((PacketSize > 1), YOU_MADE_A_PROGRAMMING_MISTAKE)
		eigen_assert(index + PacketSize - 1 < dimensions().TotalSize());

		// This code is pilfered from TensorMorphing.h
		EIGEN_ALIGN_MAX CoeffReturnType values[PacketSize];
		internal::pstore<CoeffReturnType, PacketReturnType>(values, x);
		EIGEN_UNROLL_LOOP
		for (int i = 0; i < PacketSize; ++i) {
			this->coeffRef(index + i) = values[i];
		}
	}
};

} // end namespace Eigen

#endif // EIGEN_CXX11_TENSOR_TENSOR_REVERSE_H
